
#pragma OPENCL EXTENSION cl_khr_local_int32_extended_atomics : enable
#pragma OPENCL EXTENSION cl_khr_global_int32_extended_atomics : enable

// 9. source vector is accesed as uint4 vector

__kernel void histogram( __global const uint4 *src,
         __global uint * gbuckets,
         uint min,
         uint max,
	 uint L,
	 uint nitems,
	__local * lbuckets )
{
 // 10. Set up __global memory access pattern.

 uint global_size = get_global_size(0);
 uint count = ( nitems / 4 ) / global_size;
 uint idx =  get_global_id(0);
 uint idx_saved = idx;
 uint stride =  global_size;
 uint local_idx = get_local_id(0);
 uint ws = get_local_size(0); 

 float4 minV = (float4)(min,min,min,min);
 float4 maxV = (float4)(max,max,max,max);	
// uint4 LV = uint4(L);
  	
// barrier( CLK_LOCAL_MEM_FENCE );


 for(int i = 0; i < L; i+=ws){
  idx = idx_saved;
  if ( local_idx < ws )
   lbuckets[local_idx] = 0;

  barrier( CLK_LOCAL_MEM_FENCE );
  for( int n=0; n < count; n++, idx += stride )
  {
   float4 element = (float4)((float) src[idx].x,(float)src[idx].y,(float)src[idx].z,(float)src[idx].w);
   uint4 indices = convert_uint4( ( (element - minV ) / (maxV - minV)) * L );
   if ( i <= indices.x && indices.x < i + ws ){
	(void) atom_inc(lbuckets + indices.x - i);
   }	
   if ( i <= indices.y && indices.y < i + ws ){
        (void) atom_inc(lbuckets + indices.y - i);
   } 
   if ( i <= indices.z && indices.z < i + ws ){
        (void) atom_inc(lbuckets + indices.z - i);
   } 
   if ( i <= indices.w && indices.w < i + ws ){
        (void) atom_inc(lbuckets + indices.w - i);
   } 
  }

  
  barrier( CLK_LOCAL_MEM_FENCE );

  if ( local_idx < ws )
   (void) atom_add(gbuckets + i + local_idx, lbuckets[local_idx]);
  
  barrier( CLK_GLOBAL_MEM_FENCE );

 }

// uint idx = get_global_id(0);
// uint element = src[idx];
// uint index = ((element - min)/ (max - min)) * L;
// (void ) atom_inc( gbuckets + index );
}



